__kernel void writeOnce(__global float* in, __global float* out, int stride, int N){
	float v[8];
	int id = get_global_id(0);
	
	v[0] = in[id];
	v[1] = in[id + stride];
	v[2] = in[id + stride * 2];
	v[3] = in[id + stride * 3];
	v[4] = in[id + stride * 4];
	v[5] = in[id + stride * 5];
	v[6] = in[id + stride * 6];
	v[7] = in[id + stride * 7];
	
	out[id] = v[0];
	out[id + stride] = v[1];
	out[id + stride * 2] = v[2];
	out[id + stride * 3] = v[3];
	out[id + stride * 4] = v[4];
	out[id + stride * 5] = v[5];
	out[id + stride * 6] = v[6];
	out[id + stride * 7] = v[7];
}

__kernel void writeTwice(__global float* in, __global float* out, int stride, int N){
	float v[8];
	int id = get_global_id(0);
	
	v[0] = in[id];
	v[1] = in[id + stride];
	v[2] = in[id + stride * 2];
	v[3] = in[id + stride * 3];
	v[4] = in[id + stride * 4];
	v[5] = in[id + stride * 5];
	v[6] = in[id + stride * 6];
	v[7] = in[id + stride * 7];
	
	out[id] = v[0];
	out[id + stride] = v[1];
	out[id + stride * 2] = v[2];
	out[id + stride * 3] = v[3];
	out[id + stride * 4] = v[4];
	out[id + stride * 5] = v[5];
	out[id + stride * 6] = v[6];
	out[id + stride * 7] = v[7];
	
	out[N + id] = v[0];
	out[N + id + stride] = v[1];
	out[N + id + stride * 2] = v[2];
	out[N + id + stride * 3] = v[3];
	out[N + id + stride * 4] = v[4];
	out[N + id + stride * 5] = v[5];
	out[N + id + stride * 6] = v[6];
	out[N + id + stride * 7] = v[7];
}